#include <vector>
#include <cmath>

#include "caffe/layers/sigmoid_cross_entropy_loss_layer.hpp"
#include "caffe/util/math_functions.hpp"

namespace caffe {


template <typename Dtype>
__global__ void SigmoidCrossEntropyLossForwardGPU(const int nthreads,
          const Dtype* input_data, const Dtype* target, Dtype* loss,
          const bool has_ignore_label_, const int ignore_label_,
          Dtype* counts, bool focal_loss, bool compensate_imbalance, Dtype alpha, Dtype gamma) {
  CUDA_KERNEL_LOOP(i, nthreads) {
    const int target_value = static_cast<int>(target[i]);
    if (has_ignore_label_ && target_value == ignore_label_) {
      loss[i] = 0;
      counts[i] = 0;
    } else {

      if(!focal_loss&&!compensate_imbalance)
        loss[i] = input_data[i] * (target[i] - (input_data[i] >= 0)) -
            log(1 + exp(input_data[i] - 2 * input_data[i] * (input_data[i] >= 0)));

      if(!focal_loss&&compensate_imbalance)
        loss[i] = (input_data[i] * (target[i] - (input_data[i] >= 0)) -
            log(1 + exp(input_data[i] - 2 * input_data[i] * (input_data[i] >= 0)))) * 
            (target[i]==0?(alpha):(1-alpha));

     if(focal_loss&&!compensate_imbalance)
        loss[i] = (input_data[i] * (target[i] - (input_data[i] >= 0)) -
            log(1 + exp(input_data[i] - 2 * input_data[i] * (input_data[i] >= 0)))) / 
            pow(1+exp((2*target[i]-1)*input_data[i]),gamma);

     if(focal_loss&&compensate_imbalance)
        loss[i] = input_data[i] * (target[i] - (input_data[i] >= 0)) -
          log(1 + exp(input_data[i] - 2 * input_data[i] * (input_data[i] >= 0))) * 
          (target[i]==0?(alpha):(1-alpha)) / 
          pow(1+exp((2*target[i]-1)*input_data[i]),gamma);

     counts[i] = 1;
    }
  }
}

template <typename Dtype>
__global__ void SigmoidCrossEntropyLossBackwardGPU(const int nthreads,
            const Dtype* sigmoid_output_data, const Dtype* target, Dtype* bottom_diff,
            bool focal_loss, bool compensate_imbalance, Dtype alpha, Dtype gamma){
  CUDA_KERNEL_LOOP(i,nthreads)
  {
    if(compensate_imbalance&&!focal_loss)
    {
          bottom_diff[i] = (sigmoid_output_data[i] - target[i])*((1-2*alpha)*target[i]+alpha);
    }

    if(!compensate_imbalance&&focal_loss)
    {
         if(target[i]==0)
            bottom_diff[i] = pow(sigmoid_output_data[i],gamma)*(sigmoid_output_data[i]-gamma*log(1-sigmoid_output_data[i])*(1-sigmoid_output_data[i]));
          else
            bottom_diff[i] = pow(1-sigmoid_output_data[i],gamma)*
                            (gamma*log(sigmoid_output_data[i])*sigmoid_output_data[i]-1+sigmoid_output_data[i]);
    }

    if(compensate_imbalance&&focal_loss)
    {
       if(target[i]==0)
            bottom_diff[i] = pow(sigmoid_output_data[i],gamma)*(sigmoid_output_data[i]-gamma*log(1-sigmoid_output_data[i])*(1-sigmoid_output_data[i]));
          else
            bottom_diff[i] = pow(1-sigmoid_output_data[i],gamma)*
                            (gamma*log(sigmoid_output_data[i])*sigmoid_output_data[i]-1+sigmoid_output_data[i]);

          bottom_diff[i] = bottom_diff[i]*((1-2*alpha)*target[i]+alpha);
    }
  }
}

template <typename Dtype>
__global__ void SigmoidCrossEntropyLossIgnoreDiffGPU(const int count,
    const int ignore_label, const Dtype* target, Dtype* diff) {
  CUDA_KERNEL_LOOP(i, count) {
    const int target_value = static_cast<int>(target[i]);
    if (target_value == ignore_label) {
      diff[i] = 0;
    }
  }
}

template <typename Dtype>
void SigmoidCrossEntropyLossLayer<Dtype>::Forward_gpu(
    const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
  // The forward pass computes the sigmoid outputs.
  sigmoid_bottom_vec_[0] = bottom[0];
  sigmoid_layer_->Forward(sigmoid_bottom_vec_, sigmoid_top_vec_);
  // Compute the loss (negative log likelihood)
  const int count = bottom[0]->count();
  // Stable version of loss computation from input data
  const Dtype* input_data = bottom[0]->gpu_data();
  const Dtype* target = bottom[1]->gpu_data();
  // Since this memory is not used for anything until it is overwritten
  // on the backward pass, we use it here to avoid having to allocate new GPU
  // memory to accumulate intermediate results in the kernel.
  Dtype* loss_data = bottom[0]->mutable_gpu_diff();
  Dtype* count_data = bottom[1]->mutable_gpu_diff();
  Dtype valid_count;
  if(compensate_imbalance_)
  {
    caffe_gpu_asum(bottom[1]->count(),target,&alpha_);
    alpha_ = alpha_/(bottom[1]->count());
  }
  // NOLINT_NEXT_LINE(whitespace/operators)
  SigmoidCrossEntropyLossForwardGPU<Dtype><<<CAFFE_GET_BLOCKS(count),
      CAFFE_CUDA_NUM_THREADS>>>(count, input_data, target, loss_data,
      has_ignore_label_, ignore_label_, count_data, focal_loss_, compensate_imbalance_, alpha_, gamma_);
  // Only launch another CUDA kernel if we actually need the valid count.
  if (normalization_ == LossParameter_NormalizationMode_VALID &&
      has_ignore_label_) {
    caffe_gpu_asum(count, count_data, &valid_count);
  } else {
    valid_count = count;
  }
  Dtype loss;
  caffe_gpu_asum(count, loss_data, &loss);
  normalizer_ = get_normalizer(normalization_, valid_count);
  top[0]->mutable_cpu_data()[0] = loss / normalizer_;
}

template <typename Dtype>
void SigmoidCrossEntropyLossLayer<Dtype>::Backward_gpu(
    const vector<Blob<Dtype>*>& top, const vector<bool>& propagate_down,
    const vector<Blob<Dtype>*>& bottom) {
  if (propagate_down[1]) {
    LOG(FATAL) << this->type()
               << " Layer cannot backpropagate to label inputs.";
  }
  if (propagate_down[0]) {
    // First, compute the diff
    const int count = bottom[0]->count();
    const Dtype* sigmoid_output_data = sigmoid_output_->gpu_data();
    const Dtype* target = bottom[1]->gpu_data();
    Dtype* bottom_diff = bottom[0]->mutable_gpu_diff();
    if(!focal_loss_&&!compensate_imbalance_)
    {
      caffe_copy(count, sigmoid_output_data, bottom_diff);
      caffe_gpu_axpy(count, Dtype(-1), target, bottom_diff);
    }
    else
    {
      SigmoidCrossEntropyLossBackwardGPU<Dtype><<<CAFFE_GET_BLOCKS(count),
        CAFFE_CUDA_NUM_THREADS>>>(count, sigmoid_output_data ,target, bottom_diff,
                   focal_loss_, compensate_imbalance_,alpha_, gamma_);
    }
    // Zero out gradient of ignored targets.
    if (has_ignore_label_) {
      // NOLINT_NEXT_LINE(whitespace/operators)
      SigmoidCrossEntropyLossIgnoreDiffGPU<Dtype><<<CAFFE_GET_BLOCKS(count),
        CAFFE_CUDA_NUM_THREADS>>>(count, ignore_label_, target, bottom_diff);
    }
    // Scale down gradient
    Dtype loss_weight = top[0]->cpu_diff()[0] / normalizer_;
    caffe_gpu_scal(count, loss_weight, bottom_diff);
  }
}

INSTANTIATE_LAYER_GPU_FUNCS(SigmoidCrossEntropyLossLayer);

}  // namespace caffe
